#pragma OPENCL EXTENSION cl_arm_printf : enable
__kernel void matmul_fp32(__global float *out_buf, __global float4 *lhs_buf,
                          __global float4 *rhs_buf, int4 hwk_) {
    int height = hwk_.x, width = hwk_.y, kslice = hwk_.z/4;
    int h = get_global_id(0);
    int w = get_global_id(1);    

    // 684 ms
    // float4 acc = 0;
    // for(int i=0; i<kslice; i+=4) {
    //     acc = mad(lhs_buf[h*kslice + i + 0], rhs_buf[w*kslice + i + 0], acc);
    //     acc = mad(lhs_buf[h*kslice + i + 1], rhs_buf[w*kslice + i + 1], acc);
    //     acc = mad(lhs_buf[h*kslice + i + 2], rhs_buf[w*kslice + i + 2], acc);
    //     acc = mad(lhs_buf[h*kslice + i + 3], rhs_buf[w*kslice + i + 3], acc);
    // }
    // out_buf[h*width + w] = acc.x + acc.y + acc.z + acc.w;

    // 684 ms
    // float acc = 0;
    // for(int i=0; i<kslice; i+=4) {
    //     acc += dot(lhs_buf[h*kslice + i + 0], rhs_buf[w*kslice + i + 0]);
    //     acc += dot(lhs_buf[h*kslice + i + 1], rhs_buf[w*kslice + i + 1]);
    //     acc += dot(lhs_buf[h*kslice + i + 2], rhs_buf[w*kslice + i + 2]);
    //     acc += dot(lhs_buf[h*kslice + i + 3], rhs_buf[w*kslice + i + 3]);
    // }
    // out_buf[h*width + w] = acc;
}